<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<title>Mali OpenCL SDK v1.1.0: Image Scaling</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="navtree.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="resize.js"></script>
<script type="text/javascript" src="navtree.js"></script>
<script type="text/javascript">
  $(document).ready(initResizable);
</script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td id="projectlogo"><img alt="Logo" src="logo.png"/></td>
  <td style="padding-left: 0.5em;">
   <div id="projectname">Mali OpenCL SDK v1.1.0
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.2 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Home</span></a></li>
      <li class="current"><a href="pages.html"><span>Help&#160;and&#160;Tutorials</span></a></li>
      <li><a href="files.html"><span>Files</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
</div><!-- top -->
<div id="side-nav" class="ui-resizable side-nav-resizable">
  <div id="nav-tree">
    <div id="nav-tree-contents">
      <div id="nav-sync" class="sync"></div>
    </div>
  </div>
  <div id="splitbar" style="-moz-user-select:none;" 
       class="ui-resizable-handle">
  </div>
</div>
<script type="text/javascript">
$(document).ready(function(){initNavTree('image_scaling_tutorial.html','');});
</script>
<div id="doc-content">
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Macros</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Pages</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

<div class="header">
  <div class="headertitle">
<div class="title">Image Scaling </div>  </div>
</div><!--header-->
<div class="contents">
<div class="textblock"><p>How to resize an image using OpenCL image objects.</p>
<h1><a class="anchor" id="imageScalingIntroduction"></a>
Bilinear Filtering</h1>
<p>One of the particular benefits of OpenCL image objects is the built-in bilinear filtering.</p>
<p>When you read from an OpenCL image object, instead of picking the nearest pixel to the coordinates given, it can take an average of the four closest pixels. This is hardware accelerated on a Mali-T600s series GPU inside the Texture pipeline.</p>
<p>What this means is that the quality of scaled images can be much higher for minimal perfomance cost.</p>
<p>We will use this example to provide a walkthrough of how to use OpenCL image objects.</p>
<p>Example:</p>
<div class="image">
<img src="image_scaling_bilinear_filtering.png" alt="image_scaling_bilinear_filtering.png"/>
<div class="caption">
An example of nearest pixel (left) and bilinear filtering (right).</div></div>
 <h1><a class="anchor" id="imageScalingDifferences"></a>
Image Object and Memory Buffer Differences</h1>
<p>The tutorial assumes knowledge of OpenCL memory buffers. See <a class="el" href="memory_buffers_tutorial.html">Memory Buffers</a> for more information.</p>
<p>OpenCL image objects are used almost identically to OpenCL buffers:</p>
<ul>
<li>They both have the type cl_mem.</li>
<li>For allocation, clCreateBuffer becomes clCreateImage2D (or clCreateImage3D).</li>
<li>For mapping, clEnqueueMapBuffer becomes clEnqueueMapImage.</li>
<li>For unmapping, clEnqueueUnmapBuffer works for both memory types.</li>
</ul>
<p>The biggest differences when using image objects are:</p>
<ul>
<li>Image objects require a 'sampler' in order to read from them (see <a class="el" href="image_scaling_tutorial.html#imageScalingSamplers">Samplers</a>).</li>
<li>A kernel cannot read and write to the same image (image parameters must be marked in the kernel definition with <em>__read_only</em> or <em>__write_only</em>).</li>
<li>Images have a defined data format.</li>
</ul>
<h1><a class="anchor" id="imageScalingSamplers"></a>
Samplers</h1>
<p>As mentioned, in order to be able to read from an image object you require a sampler.</p>
<p>A sampler defines:</p>
<ul>
<li>whether the coordinates you are using are normalized<ul>
<li>normalized (in the range [0,1])</li>
<li>non-normalized</li>
</ul>
</li>
<li>what strategy to use when the coordinates are outside the bounds of the image<ul>
<li>none (you guarantee that the coordinates are inside the bounds)</li>
<li>clamp to edge (returns the color of the closest valid pixel)</li>
<li>clamp (returns a border color defined by the image format)</li>
<li>repeat (acts as if there are infinite copies of the image tiled next to each other)</li>
<li>mirrored repeat (the same as repeat except coordinates are flipped at every edge)</li>
</ul>
</li>
<li>which filtering strategy to use<ul>
<li>nearest</li>
<li>bilinear</li>
</ul>
</li>
</ul>
<p>Certain combinations of these options are restricted.</p>
<p>Samplers can be defined on the host side using clCreateSampler() and passed in the kernel as an argument or directly inside the kernel source. Passing the sampler as an argument to the kernel gives you the flexibility to run the same kernel with different sampling options.</p>
<h1><a class="anchor" id="imageScalingCode"></a>
Resizing images with bilinear filtering</h1>
<p>Unless otherwise noted, all code snippets come from <a class="el" href="image__scaling_8cpp.html">image_scaling.cpp</a>.</p>
<p>In the sample code we will resize an input image using OpenCL. The image is scaled up by a factor of 8 (adjustable in the code) with bilinear filtering enabled.</p>
<ol type="1">
<li><p class="startli"><b>Allocate memory for your images</b></p>
<p class="startli">Image objects are allocated in almost the same way as buffers. The main difference is having to specify the data format of the image being used.</p>
<p class="startli">You can list the avaliable image formats on your platform using the <a class="el" href="common_8h.html#a736045ae74d9ed8d35cacce59b92772a">printSupported2DImageFormats</a> method.</p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Specify the format of the image.</span></div>
<div class="line"><span class="comment">     * The bitmap image we are using is RGB888, which is not a supported OpenCL image format.</span></div>
<div class="line"><span class="comment">     * We will use RGBA8888 and add an empty alpha channel.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    cl_image_format format;</div>
<div class="line">    format.image_channel_data_type = CL_UNORM_INT8;</div>
<div class="line">    format.image_channel_order = CL_RGBA;</div>
<div class="line"></div>
<div class="line">    <span class="comment">/* Allocate memory for the input image that can be accessed by the CPU and GPU. */</span></div>
<div class="line">    <span class="keywordtype">bool</span> createMemoryObjectsSuccess = <span class="keyword">true</span>;</div>
<div class="line"></div>
<div class="line">    memoryObjects[0] = clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, &amp;format, width, height, 0, NULL, &amp;errorNumber);</div>
<div class="line">    createMemoryObjectsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(errorNumber);</div>
<div class="line"></div>
<div class="line">    memoryObjects[1] = clCreateImage2D(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, &amp;format, newWidth, newHeight, 0, NULL, &amp;errorNumber);</div>
<div class="line">    createMemoryObjectsSuccess &amp;= <a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(errorNumber);</div>
<div class="line"></div>
<div class="line">    <span class="keywordflow">if</span> (!createMemoryObjectsSuccess)</div>
<div class="line">    {</div>
<div class="line">        <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);</div>
<div class="line">        cerr &lt;&lt; <span class="stringliteral">&quot;Failed creating the image. &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">        <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Map the memory to a host pointer</b></p>
<p class="startli">Again, this step is very similar to mapping a buffer:</p>
<div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Like with memory buffers, we now map the allocated memory to a host side pointer.</span></div>
<div class="line"><span class="comment">     * Unlike buffers, we must specify origin coordinates, width and height for the region of the image we wish to map.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">size_t</span> origin[3] = {0, 0, 0};</div>
<div class="line">    <span class="keywordtype">size_t</span> region[3] = {width, height, 1};</div>
<div class="line"></div>
<div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * clEnqueueMapImage also returns the rowPitch; the width of the mapped region in bytes.</span></div>
<div class="line"><span class="comment">     * If the image format is not known, this is required information when accessing the image object as a normal array.</span></div>
<div class="line"><span class="comment">     * The number of bytes per pixel can vary with the image format being used,</span></div>
<div class="line"><span class="comment">     * this affects the offset into the array for a given coordinate.</span></div>
<div class="line"><span class="comment">     * In our case the image format is fixed as RGBA8888 so we don&#39;t need to worry about the rowPitch.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    <span class="keywordtype">size_t</span> rowPitch;</div>
<div class="line"></div>
<div class="line">    <span class="keywordtype">unsigned</span> <span class="keywordtype">char</span>* inputImageRGBA = (<span class="keywordtype">unsigned</span> <span class="keywordtype">char</span>*)clEnqueueMapImage(commandQueue,  memoryObjects[0], CL_TRUE, CL_MAP_WRITE, origin, region, &amp;rowPitch, NULL, 0, NULL, NULL, &amp;errorNumber);</div>
<div class="line">    <span class="keywordflow">if</span> (!<a class="code" href="common_8cpp.html#a73727ed6f5d821d8c97756404e145644" title="Check an OpenCL error number for errors.">checkSuccess</a>(errorNumber))</div>
<div class="line">    {</div>
<div class="line">        <a class="code" href="common_8cpp.html#acd459e74d0cef3c4616c7ce1a5a47f4d" title="Release any OpenCL objects that have been created.">cleanUpOpenCL</a>(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);</div>
<div class="line">        cerr &lt;&lt; <span class="stringliteral">&quot;Failed mapping the input image. &quot;</span> &lt;&lt; __FILE__ &lt;&lt; <span class="stringliteral">&quot;:&quot;</span>&lt;&lt; __LINE__ &lt;&lt; endl;</div>
<div class="line">        <span class="keywordflow">return</span> 1;</div>
<div class="line">    }</div>
</div><!-- fragment --></li>
<li><p class="startli"><b>Initialize the memory</b></p>
<p class="startli">Use the pointer on the host side to fill the image with data.</p>
</li>
<li><p class="startli"><b>Unmap the memory</b></p>
<p class="startli">Unmap the pointer on the host side (using clEnqueueUnmapBuffer just like with buffers) so that the data can be used inside the kernel.</p>
</li>
<li><p class="startli"><b>Pass the image to the kernel</b></p>
<p class="startli">Pass the images to the kernel as arguments in the same way as buffers.</p>
</li>
<li><p class="startli"><b>Use the image inside the kernel</b></p>
<p class="startli">Code snippets in this section are from <a class="el" href="image__scaling_8cl.html">image_scaling.cl</a>.</p>
<ol type="a">
<li><p class="startli"><b>Define a sampler</b></p>
<p class="startli">See <a class="el" href="image_scaling_tutorial.html#imageScalingSamplers">Samplers</a> for more information.</p>
<p class="startli">In this walkthrough we define the sample in the kernel source:</p>
<div class="fragment"><div class="line"><span class="keyword">const</span> sampler_t <a class="code" href="image__scaling_8cl.html#a9165c6046146d2d28cc268394d585f17">sampler</a> = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR;</div>
</div><!-- fragment --></li>
<li><b>Calculate the coordinates</b> <div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * There is one kernel instance per pixel in the destination image.</span></div>
<div class="line"><span class="comment">     * The global id of this kernel instance is therefore a coordinate in the destination image.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    int2 coordinate = (int2)(get_global_id(0), get_global_id(1));</div>
<div class="line"></div>
<div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * That coordinate is only valid for the destination image.</span></div>
<div class="line"><span class="comment">     * If we normalize the coordinates to the range [0.0, 1.0] (using the height and width of the destination image),</span></div>
<div class="line"><span class="comment">     * we can use them as coordinates in the sourceImage.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    float2 normalizedCoordinate = convert_float2(coordinate) * (float2)(widthNormalizationFactor, heightNormalizationFactor);</div>
</div><!-- fragment --></li>
<li><b>Read the source image</b> <div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Read colours from the source image.</span></div>
<div class="line"><span class="comment">     * The sampler determines how the coordinates are interpreted.</span></div>
<div class="line"><span class="comment">     * Because bilinear filtering is enabled, the value of colour will be the average of the 4 pixels closest to the coordinate.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    float4 colour = read_imagef(sourceImage, sampler, normalizedCoordinate);</div>
</div><!-- fragment --></li>
<li><b>Write the destination image</b> <div class="fragment"><div class="line">    <span class="comment">/*</span></div>
<div class="line"><span class="comment">     * Write the colour to the destination image.</span></div>
<div class="line"><span class="comment">     * No sampler is used here as all writes must specify an exact valid pixel coordinate.</span></div>
<div class="line"><span class="comment">     */</span></div>
<div class="line">    write_imagef(destinationImage, coordinate, colour);</div>
</div><!-- fragment --></li>
</ol>
</li>
<li><p class="startli"><b>Retrieve the results</b></p>
<p class="startli">Map the image object to a host pointer and read the results.</p>
</li>
</ol>
<h1><a class="anchor" id="imageScalingRunning"></a>
Running the Sample</h1>
<ol type="1">
<li><p class="startli">From a command prompt in the root of the SDK, run:</p>
<div class="fragment"><div class="line">cd samples/<a class="code" href="image__scaling_8cl.html#a2f2c7656161bcd4a3956c362d0052b26" title="Image scaling kernel function.">image_scaling</a></div>
</div><!-- fragment --> <div class="fragment"><div class="line">make install</div>
</div><!-- fragment --><p class="startli">This compiles the image scaling sample code and copies all the files it needs to run to the bin folder in the root directory of the SDK.</p>
</li>
<li>Copy this folder to the board.</li>
<li><p class="startli">Navigate to the folder on the board and run the image scaling binary:</p>
<div class="fragment"><div class="line">./<a class="code" href="image__scaling_8cl.html#a2f2c7656161bcd4a3956c362d0052b26" title="Image scaling kernel function.">image_scaling</a></div>
</div><!-- fragment --></li>
<li><p class="startli">You should see output similar to:</p>
<div class="fragment"><div class="line">11 Image formats supported (channel order, channel data type):</div>
<div class="line">CL_RGBA, CL_UNORM_INT8</div>
<div class="line">CL_RGBA, CL_UNORM_INT16</div>
<div class="line">CL_RGBA, CL_SIGNED_INT8</div>
<div class="line">CL_RGBA, CL_SIGNED_INT16</div>
<div class="line">CL_RGBA, CL_SIGNED_INT32</div>
<div class="line">CL_RGBA, CL_UNSIGNED_INT8</div>
<div class="line">CL_RGBA, CL_UNSIGNED_INT16</div>
<div class="line">CL_RGBA, CL_UNSIGNED_INT32</div>
<div class="line">CL_RGBA, CL_HALF_FLOAT</div>
<div class="line">CL_RGBA, CL_FLOAT</div>
<div class="line">CL_BGRA, CL_UNORM_INT8</div>
<div class="line"></div>
<div class="line">Profiling information:</div>
<div class="line">Queued time:    0.092ms</div>
<div class="line">Wait time:      0.135206ms</div>
<div class="line">Run time:       31.5405ms</div>
</div><!-- fragment --><p class="startli">An output image should be created on the board called output.bmp.</p>
</li>
</ol>
<p>Find solutions for <a class="el" href="common_issues.html">Common Issues</a>.</p>
<h1><a class="anchor" id="imageScalingMoreInformation"></a>
More Information</h1>
<p>For more information have a look at the code in <a class="el" href="image__scaling_8cpp.html">image_scaling.cpp</a> and <a class="el" href="image__scaling_8cl.html">image_scaling.cl</a>. </p>
</div></div><!-- contents -->
</div><!-- doc-content -->
<!-- start footer part -->
<div id="nav-path" class="navpath"><!-- id is needed for treeview function! -->
  <ul>
    <li class="navelem"><a class="el" href="tutorials.html">Tutorials</a></li><li class="navelem"><a class="el" href="image_objects_tutorial.html">Image Objects</a></li>
    <li class="footer">
        <a href="http://www.arm.com/">(C) ARM Ltd. 2013</a>
    </li>
  </ul>
</div>
</body>
</html>
